#pragma OPENCL EXTENSION cl_khr_subgroups : enable

__kernel void subgroup_test(__global int* output) {
    uint gid = get_global_id(0);
    uint lid = get_local_id(0);
    uint subgroup_id = get_sub_group_id();
    uint subgroup_size = get_sub_group_size();
    uint local_subid = get_sub_group_local_id();

    // 使用sub_group_broadcast()验证功能
    uint broadcast_val = sub_group_broadcast(lid, 0);

    // 把结果写入输出数组（混合信息编码）
    output[gid] = (subgroup_id << 24) | (subgroup_size << 16) | (local_subid << 8) | (broadcast_val & 0xFF);
}